//#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 220
#include "quant_matmul_v2_pertoken_basic.h"
#include "kernel_operator.h"


using namespace AscendC;
using namespace matmul;

extern "C" __global__ __aicore__ void quant_matmul_v2(GM_ADDR x1, GM_ADDR x2, GM_ADDR scale, GM_ADDR bias, 
                                                    GM_ADDR pertoken_scale, GM_ADDR y, 
                                                    GM_ADDR buff0,
                                                    GM_ADDR buff1,
                                                    GM_ADDR buff2,
                                                    GM_ADDR buff3,
                                                    GM_ADDR buff4,
                                                    GM_ADDR buff5,
                                                    GM_ADDR buff6,
                                                    GM_ADDR buff7,
                                                    GM_ADDR workspace,
                                                    GM_ADDR tiling) {
    if (workspace == nullptr){
        return;
    }
    
    GM_ADDR user1 = GetUserWorkspace(workspace);
    if(user1 == nullptr){
        return;
    }
    const bool transpose_x1 = false;
    const bool transpose_x2 = false;
    GET_TILING_DATA(tiling_data, tiling);
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
    const QuantMatmulV2TilingData *qmm_v2_tiling_data = &tiling_data;
    if (TILING_KEY_IS(10100)) {
    TPipe t_pipe;
    BmmDequantPertokenBasic<DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_Y, FORMAT_X1, FORMAT_X2, transpose_x1, 
                            transpose_x2, QuantMatmulV2Update> op;
    op.init(x1, x2, scale, bias, pertoken_scale, y, 
    buff0, buff1, buff2, buff3, buff4, buff5, buff6, buff7, 
    user1, qmm_v2_tiling_data, &t_pipe);
    // op.process_event();
    op.process();
    t_pipe.Destroy();}
}


    